[copilot mirror] B1: TopologyBuilder resource seam + cuda::Buffer naming by harrism · Pull Request #1 · harrism/openvdb

harrism · 2026-08-05T04:41:14Z

Copilot-review mirror of AcademySoftwareFoundation#2268 (fork-side, so Copilot can be requested). Fixes land on the shared branch and flow to the upstream PR.

Name the free operation destroy, as cuda::buffer does, and add the stream-taking overload it also provides. clear stays but only as a transitional delegate, marked as such: it exists because GridHandle::reset still calls it, and goes away with the legacy dual buffers that cuda::Buffer replaces. Rename setStream to set_stream. Member names cannot be aliased, so matching the standard spelling is the whole reason the resource concept kept allocate_async; the same argument applies here. Note in passing that ours deliberately does not synchronize, matching cuda::buffer's set_stream_unsynchronized rather than its set_stream, whose own documentation and implementation disagree (NVIDIA/cccl#10649). Add swap. The generic std::swap already does the right thing through the move operations, but both std::vector and rmm::device_buffer provide one. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>

Eleven of the builder's buffers are device-only -- nothing reads them on the host -- yet they were cuda::DeviceBuffer, whose host pointer and per-device array they never use. Move them to the single-space cuda::Buffer and give the builder a resource parameter, so its scratch is injectable like PointsToGrid's already is, and freed by scope rather than by hand. mProcessedRoot and mData are read on the host and stay dual. This is not a speedup: DeviceBuffer::init allocates host or device memory but never both, and these buffers always passed a real device id, so no cudaMallocHost was on this path to begin with. Measured dilate on 1k/20k/ 200k points, and the difference is within run-to-run noise. Hoist the Data struct out of the class. It does not depend on the resource, and leaving it nested would give every ResourceT its own incompatible type for the device functors to name. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>

Copilot

Pull request overview

Warning

Copilot couldn't run its full agentic review because it didn't start before the timeout. Make sure your repository has a runner available, or add a copilot-code-review.yml file specifying one with the runs-on attribute. See the docs for more details.

This PR mirrors upstream work to introduce a resource-injectable CUDA scratch allocation seam in tools::cuda::TopologyBuilder, while aligning NanoVDB CUDA buffer APIs with cuda::buffer naming (destroy, set_stream, swap) and updating call sites/tests accordingly.

Changes:

Added nanovdb::cuda::Buffer / BufferView changelog entry and extended Buffer with destroy(stream) and swap, plus set_stream naming.
Refactored tools::cuda::TopologyBuilder to be templated on a CUDA memory resource and to use nanovdb::cuda::Buffer for device-only scratch allocations.
Updated CUDA tool kernels and unit tests to use the new names and pointer accessors.

Reviewed changes

Copilot reviewed 7 out of 7 changed files in this pull request and generated 3 comments.

Show a summary per file

File	Description
pendingchanges/nanovdb.txt	Documents the new CUDA buffer types and resource concepts.
nanovdb/nanovdb/unittest/TestBuffer.cu	Adds/updates unit tests for `destroy`, `destroy(stream)`, `swap`, and `set_stream`.
nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh	Introduces resource-injected scratch allocations via `cuda::Buffer` and factors out shared `TopologyBuilderData`.
nanovdb/nanovdb/tools/cuda/RefineGrid.cuh	Updates kernel argument pointers to use `.data()` on new scratch buffers.
nanovdb/nanovdb/tools/cuda/PruneGrid.cuh	Updates kernel argument pointers to use `.data()` on new scratch buffers.
nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh	Updates kernel argument pointers to use `.data()` on new scratch buffers.
nanovdb/nanovdb/cuda/Buffer.h	Renames to `set_stream`, introduces `destroy`, `destroy(stream)`, `swap`, and keeps `clear()` as transitional alias.

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

harrism · 2026-08-05T05:08:05Z

+    ScratchT upperCountsBuffer = ScratchT(stream, *mResource, processedTileCount*sizeof(uint32_t), nanovdb::cuda::noInit);
+    ScratchT lowerCountsBuffer = ScratchT(stream, *mResource, size*sizeof(uint32_t), nanovdb::cuda::noInit);
+    ScratchT leafCountsBuffer = ScratchT(stream, *mResource, size*sizeof(uint32_t), nanovdb::cuda::noInit);

    using CountType = uint32_t (*)[Mask<5>::SIZE];
-    auto lowerCounts = reinterpret_cast<CountType>( lowerCountsBuffer.deviceData() );
-    auto leafCounts = reinterpret_cast<CountType>( leafCountsBuffer.deviceData() );
+    auto lowerCounts = reinterpret_cast<CountType>(lowerCountsBuffer.data());
+    auto leafCounts = reinterpret_cast<CountType>(leafCountsBuffer.data());


Not applicable: Buffer does not derive alignment from T — allocate passes R::DEFAULT_ALIGNMENT (256 for DeviceResource, and cudaMallocAsync 256-aligns regardless), the same alignment the raw allocate_async calls used before this change. The reinterpret_casts are to types with far smaller alignment requirements. Typed buffers are a possible future polish, but there is no misalignment here.

harrism · 2026-08-05T05:08:07Z

    mProcessedRoot.clear(stream);
-    mUpperMasks.clear(stream);
-    mLowerMasks.clear(stream);
-    mLowerOffsets.clear(stream);
-    mLeafOffsets.clear(stream);
-}// TopologyBuilder<BuildT>::processLowerNodes
+    mUpperMasks.destroy();
+    mLowerMasks.destroy();
+    mLowerOffsets.destroy();
+    mLeafOffsets.destroy();
+}// TopologyBuilder<BuildT, ResourceT>::processLowerNodes


Fixed in c7302bf — all seven release sites now pass the caller's stream to destroy(stream). The retained stream and the caller's stream are identical on every current path (the builder is constructed with the operator's stream and every call uses it), but the explicit form no longer depends on that invariant.

harrism · 2026-08-05T05:08:08Z

    int device = 0;
    cudaGetDevice(&device);
-    nanovdb::cuda::DeviceBuffer upperCountsBuffer = nanovdb::cuda::DeviceBuffer::create(processedTileCount*sizeof(uint32_t), nullptr, device, stream);
-    nanovdb::cuda::DeviceBuffer lowerCountsBuffer = nanovdb::cuda::DeviceBuffer::create(size*sizeof(uint32_t), nullptr, device, stream);
-    nanovdb::cuda::DeviceBuffer leafCountsBuffer = nanovdb::cuda::DeviceBuffer::create(size*sizeof(uint32_t), nullptr, device, stream);
+    ScratchT upperCountsBuffer = ScratchT(stream, *mResource, processedTileCount*sizeof(uint32_t), nanovdb::cuda::noInit);


Fixed in c7302bf — removed the four cudaGetDevice pairs left unused by the conversion (allocateInternalMaskBuffers, countNodes, processLowerNodes, processLeafOffsets). getBuffer keeps its copy, which still feeds BufferT::create.

Each release site sits in a function that receives the stream, so pass it to destroy rather than relying on the retained stream matching -- they are the same on every current path, but the explicit form does not depend on that staying true. Also drop the cudaGetDevice calls whose result the Buffer conversion left unused. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>

Copilot

Pull request overview

Copilot reviewed 7 out of 7 changed files in this pull request and generated 7 comments.

Suppressed comments (2)

nanovdb/nanovdb/cuda/Buffer.h:281

Since clear() is intended to be transitional and 'not the name to use', consider marking it as deprecated (e.g., [[deprecated(\"use destroy()\")]]) and/or gating it behind a compatibility macro. That will help prevent new call sites from continuing to adopt clear() and makes the eventual removal safer.

    /// @brief Frees the buffer memory (if any) and resets to the empty state.
    /// @note Transitional, and not the name to use: it exists only because
    ///       GridHandle::reset still calls clear() on its buffer. It goes away
    ///       when the legacy dual buffers do and GridHandle moves to destroy().
    void clear() { this->destroy(); }

pendingchanges/nanovdb.txt:5

This release-note bullet is very long and packs multiple concepts (Buffer, BufferView, stream-ordering semantics, GridHandle wrapping, naming conventions, and new resource concept) into a single line. Consider splitting into 2–3 shorter bullets or sentences to improve readability and make it easier to scan in change logs.

  - Added nanovdb::cuda::Buffer and nanovdb::cuda::BufferView (CUDA): a typed, resource-aware, stream-ordered container that allocates from an injectable memory resource and frees on its retained stream, and a non-owning view over externally managed memory that a GridHandle can wrap without copying. Member names follow cuda::buffer (destroy, set_stream, swap). Also added the synchronous resource concept nanovdb::cuda::is_resource alongside is_async_resource.

harrism · 2026-08-05T05:30:56Z

+    /// @brief Device-only scratch storage, allocated from the injected
+    ///        resource. These buffers are never read on the host, so they use
+    ///        the single-space Buffer rather than the dual DeviceBuffer, whose
+    ///        host pointer and per-device array they would leave unused.
+    using ScratchT = nanovdb::cuda::Buffer<std::byte, ResourceT>;


Dismissed for the reason given on the earlier thread: Buffer allocates with R::DEFAULT_ALIGNMENT (256 for every CUDA-backed resource), not alignof(T), so the reinterpretation is to types with far weaker alignment requirements — identical to the raw allocate_async calls before this change. That said, the custom-resource corner is now closed at compile time: 2608ea8 (stacked B2 PR) asserts ResourceT::DEFAULT_ALIGNMENT >= alignof(uint64_t), so a resource with a weaker guarantee fails with a diagnostic instead of misaligning.

harrism · 2026-08-05T05:30:57Z

+    auto lowerCounts = reinterpret_cast<CountType>(lowerCountsBuffer.data());
+    auto leafCounts = reinterpret_cast<CountType>(leafCountsBuffer.data());


See the reply on the thread at line 63 — same finding; the alignment assert in 2608ea8 covers this site too.

harrism · 2026-08-05T05:30:58Z

+        reinterpret_cast<uint32_t*>(lowerCountsBuffer.data()),
+        reinterpret_cast<uint32_t*>(mLowerOffsets.data())+1,
        size);
-    cudaCheck(cudaMemcpyAsync(&data()->nodeCount[1], static_cast<uint32_t*>(mLowerOffsets.deviceData())+size, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream));
+    cudaCheck(cudaMemcpyAsync(&data()->nodeCount[1], reinterpret_cast<uint32_t*>(mLowerOffsets.data())+size, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream));


See the reply on the thread at line 63 — same finding; the alignment assert in 2608ea8 covers this site too.

harrism · 2026-08-05T05:30:59Z

    mProcessedRoot.clear(stream);
-    mUpperMasks.clear(stream);
-    mLowerMasks.clear(stream);
-    mLowerOffsets.clear(stream);
-    mLeafOffsets.clear(stream);
-}// TopologyBuilder<BuildT>::processLowerNodes
+    mUpperMasks.destroy(stream);
+    mLowerMasks.destroy(stream);
+    mLowerOffsets.destroy(stream);
+    mLeafOffsets.destroy(stream);
+}// TopologyBuilder<BuildT, ResourceT>::processLowerNodes


The premise is right — destroy(stream) requires an async resource — but a synchronous ResourceT has never compiled here (TempPool requires the async pair), and the stack resolves it deliberately: f619056 asserts is_async_resource<ResourceT> at class scope with a diagnostic naming the builder, and synchronous resources drive the builders through AsyncFromSync (upstream AcademySoftwareFoundation#2272), which models AsyncResource, so destroy(stream) stays well-formed. An if constexpr fork at each of the seven sites would duplicate the dispatch the wrapper centralizes.

harrism · 2026-08-05T05:31:01Z

+        data()->nodeCount[0], topology::detail::UpdateAndPropagateLeafBBoxFunctor<BuildT>(), deviceData(), reinterpret_cast<uint32_t*>(mLeafParents.data()));
+    mLeafParents.destroy(stream);


See the reply on the thread at line 476 — same finding, resolved by the class-scope assert (f619056) plus AsyncFromSync (AcademySoftwareFoundation#2272).

harrism · 2026-08-05T05:31:02Z

+        data()->nodeCount[1], topology::detail::PropagateLowerBBoxFunctor<BuildT>(), deviceData(), reinterpret_cast<uint32_t*>(mLowerParents.data()));
+    mLowerParents.destroy(stream);


See the reply on the thread at line 476 — same finding, resolved by the class-scope assert (f619056) plus AsyncFromSync (AcademySoftwareFoundation#2272).

harrism · 2026-08-05T05:31:03Z

+        util::cuda::lambdaKernel<<<1, 1, 0, stream>>>(1, topology::detail::PostProcessGridTreeFunctor<BuildT>(), deviceData(), reinterpret_cast<uint64_t*>(mVoxelOffsets.data()));
    cudaCheckError();
-    mVoxelOffsets.clear(stream);
+    mVoxelOffsets.destroy(stream);


See the reply on the thread at line 476 — same finding, resolved by the class-scope assert (f619056) plus AsyncFromSync (AcademySoftwareFoundation#2272).

harrism · 2026-08-05T05:48:50Z

Copilot review cycles converged (final cycle: no new findings). All accepted changes are on the shared branch and reflected in the upstream PR; threads here document the dismissals. Branch retained.

harrism and others added 3 commits August 5, 2026 01:18

NanoVDB: note cuda::Buffer and cuda::BufferView in pendingchanges

fb05c71

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>

harrism requested a lite review from Copilot August 5, 2026 04:41

Copilot AI reviewed Aug 5, 2026

View reviewed changes

Copilot started reviewing on behalf of harrism August 5, 2026 05:07 View session

harrism requested a lite review from Copilot August 5, 2026 05:08

Copilot AI reviewed Aug 5, 2026

View reviewed changes

Copilot started reviewing on behalf of harrism August 5, 2026 05:30 View session

harrism mentioned this pull request Aug 5, 2026

NanoVDB: inject a resource into TopologyBuilder scratch, and align cuda::Buffer with cuda::buffer (CUDA) AcademySoftwareFoundation/openvdb#2268

Merged

harrism closed this Aug 5, 2026

Navigation Menu

[copilot mirror] B1: TopologyBuilder resource seam + cuda::Buffer naming - #1

[copilot mirror] B1: TopologyBuilder resource seam + cuda::Buffer naming#1
harrism wants to merge 4 commits into
masterfrom
nanovdb-buffer-retrofit

harrism commented Aug 5, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

harrism Aug 5, 2026

Uh oh!

harrism Aug 5, 2026

Uh oh!

harrism Aug 5, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

harrism Aug 5, 2026

Uh oh!

harrism Aug 5, 2026

Uh oh!

harrism Aug 5, 2026

Uh oh!

harrism Aug 5, 2026

Uh oh!

harrism Aug 5, 2026

Uh oh!

harrism Aug 5, 2026

Uh oh!

harrism Aug 5, 2026

Uh oh!

harrism commented Aug 5, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

		auto lowerCounts = reinterpret_cast<CountType>(lowerCountsBuffer.data());
		auto leafCounts = reinterpret_cast<CountType>(leafCountsBuffer.data());

		data()->nodeCount[0], topology::detail::UpdateAndPropagateLeafBBoxFunctor<BuildT>(), deviceData(), reinterpret_cast<uint32_t*>(mLeafParents.data()));
		mLeafParents.destroy(stream);

		data()->nodeCount[1], topology::detail::PropagateLowerBBoxFunctor<BuildT>(), deviceData(), reinterpret_cast<uint32_t*>(mLowerParents.data()));
		mLowerParents.destroy(stream);

Conversation

harrism commented Aug 5, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

harrism commented Aug 5, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants